/*
 * Projekt: GMU 2013 - Canny hranovy detektor OpenCL/CUDA
 * Autori: Pavel Nevrala - xnevra00 (xnevra00@stud.fit.vutbr.cz)
 *         Martin Jurak - xjurak02 (xjurak02@stud.fit.vutbr.cz)
 */
#include "DoubleThresholdingCL.h"
#include <cmath>

DoubleThresholdingCL::DoubleThresholdingCL(cl_uint maxWorkGroupP, cl_context *contextP, cl_command_queue *commandQueueP, cl_program *programP, cl_float *edgesDirectionP):
	CannyPartsBase(maxWorkGroupP), context(contextP), commandQueue(commandQueueP), program(programP), edgesDirection(edgesDirectionP)
{
}


DoubleThresholdingCL::~DoubleThresholdingCL(void)
{
}

/*
 * Provede nastaveni OpenCl a kernelu pro doublethresholding s hysterezi.
 * @image - vstupni obraz
 * @width - sirka obrazu
 * @height - vyska obrazu	 
 */
double DoubleThresholdingCL::run(cl_float *image, cl_uint width, cl_uint height)
{
	return run(image, width,height,DEFAULT_THR_LOW, DEFAULT_THR_HIGH);
}

/*
 * Provede nastaveni OpenCl a kernelu pro doublethresholding s hysterezi s nastavenim prahu.
 * @image - vstupni obraz
 * @width - sirka obrazu
 * @height - vyska obrazu	
 * @thrLow - dolni prah
 * @thrHigh - horni prah
 * @return - doba vypoctu na GPU
 */
double DoubleThresholdingCL::run(cl_float *image, cl_uint width, cl_uint height, cl_uchar thrLow, cl_uchar thrHigh)
{
	cl_int clErr;

	//vytvoreni bufferu
	cl_mem cl_in = clCreateBuffer(*context, CL_MEM_READ_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_in, CL_TRUE /*blocking write*/,	0, width * height * sizeof(cl_float), image, 0, NULL, NULL);
	cl_mem cl_edgesDirection = clCreateBuffer(*context, CL_MEM_READ_ONLY, width * height * sizeof(cl_float), NULL, &clErr);
	clErr = clEnqueueWriteBuffer(*commandQueue, cl_edgesDirection, CL_TRUE /*blocking write*/,	0, width * height * sizeof(cl_float), edgesDirection, 0, NULL, NULL);
	cl_mem cl_out = clCreateBuffer(*context, CL_MEM_WRITE_ONLY, width * height * sizeof(cl_float), NULL, &clErr);

	//ziskani kernelu
	cl_kernel kernelDoubleThresholding = clCreateKernel(*program, "doubleThresholding", &clErr);
	errorController.CheckOpenCLError( clErr, "clCreateKernel DoubleThresholding" );	
	
// ========================================================================================================================================================================================
// ========================================================================================================================================================================================

   	clErr = clSetKernelArg(kernelDoubleThresholding, 0, sizeof(cl_mem), &cl_in);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImageDoubleThr)");
	clErr = clSetKernelArg(kernelDoubleThresholding, 1, sizeof(cl_mem), &cl_edgesDirection);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (inputImageDoubleThr)");
	clErr = clSetKernelArg(kernelDoubleThresholding, 2, sizeof(cl_mem), &cl_out);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImageDoubleThr)");	
	clErr = clSetKernelArg(kernelDoubleThresholding, 3, sizeof(cl_uchar), &thrLow);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImageDoubleThr)");	
	clErr = clSetKernelArg(kernelDoubleThresholding, 4, sizeof(cl_uchar), &thrHigh);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (outImageDoubleThr)");	
	clErr = clSetKernelArg(kernelDoubleThresholding, 5, sizeof(cl_uint), &width);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (widthImage)");
	clErr = clSetKernelArg(kernelDoubleThresholding, 6, sizeof(cl_uint), &height);
	errorController.CheckOpenCLError(clErr, "clSetKernelArg. (heightImage)");

	size_t GWS[2], LWS[2];
	// nastaveni velikosti work-group
	int localGroup = (int)sqrt((double)maxWorkGroup);
	GWS[0]=((width/localGroup)+1)*localGroup;
	GWS[1]=((height/localGroup)+1)*localGroup;
	LWS[0]=localGroup;
	LWS[1]=localGroup;
		
	cl_event event_DoubleThr;
	clErr = clEnqueueNDRangeKernel(*commandQueue, kernelDoubleThresholding, 2, NULL, GWS, LWS, 0, NULL, &event_DoubleThr);
	errorController.CheckOpenCLError(clErr, "enquekernelDoubleThresholding");
		
	clErr = clWaitForEvents(1, &event_DoubleThr);	
	errorController.CheckOpenCLError(clErr, "waitforevents DoubleThr");

	// zjisteni doby vypoctu
	clGetEventProfilingInfo(event_DoubleThr, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
	clGetEventProfilingInfo(event_DoubleThr, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);
	total_time = (double)(time_end - time_start);

// =======================================================================================================================================================================================
// ========================================================================================================================================================================================

	// zapis zpracovaneho obrazu
	outImage = new cl_float[width * height];
	memset(outImage, 0, width * height * sizeof(cl_float));
	clErr = clEnqueueReadBuffer (*commandQueue, cl_out, CL_TRUE, 0, width * height * sizeof(cl_float), outImage, 0, NULL, NULL);
	errorController.CheckOpenCLError(clErr, "waitforevents DoubleThr");

	// Uvolneni zdroju
    clErr = clReleaseKernel(kernelDoubleThresholding);
    errorController.CheckOpenCLError(clErr, "clReleaseKernel DoubleThr.");
	clErr = clReleaseMemObject(cl_in);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject input");     
    clErr = clReleaseMemObject(cl_out);
    errorController.CheckOpenCLError(clErr, "clReleaseMemObject out");	

	// doba vypoctu v milisekundach
	return total_time / 1000000.0;
    
}